Repository navigation
fix(shader): gather a 3D image's image_gather4_lz from one slice as measured - #9183
Conversation
|
Some comparisons could not be completed because a pull request conflicts with the base branch. Review these shared files manually before concluding that there is no overlap.
|
jotitadev
left a comment
There was a problem hiding this comment.
First Apple GPU / MoltenVK run. It confirms the bug here as a distinct vendor failure, shows the fix clears it, and surfaces one case in the new test that cannot run on this device.
Machine: Apple M2 Max, macOS 27.2, x86_64 under Rosetta 2. Vulkan loader + MoltenVK 1.4.2 (Vulkan SDK 1.4.363.0). Release, BUILD_TESTING=ON.
Base/method: PR #8658 feat/macos-port @ a1fb8e3e9c (current main a6dcea9279 plus the macOS port), not main — main does not build on macOS. This PR's diff (git diff merge-base..df4fd534d -- . ':!docs') applied on top, clean.
Negative control (revert the production files to the PR's base, keep the new ImageGatherVolume.cpp): agc_driver_image_gather_volume fails at pipeline creation. MoltenVK translates the OpImageGather on a 3D image to MSL and Metal refuses it:
program_source:84:49: error: no member named 'gather' in 'metal::texture3d<float, metal::access::sample, metal::memory_coherence_threadgroup>'
float4 _544 = spvDescriptorSet0.image_7[0u].gather(spvDescriptorSet0.samplers[0u], float3(...), component::x);
vkCreateComputePipelines: Vulkan result -3
So this is a fourth vendor's manifestation of the same bug the PR fixes: alongside the RTX 5080 SIGILL, the RADV crash and lavapipe's -13, MoltenVK rejects texture3d.gather at MSL compile with VK_ERROR_INITIALIZATION_FAILED.
With this PR: the shader compiles (the fix emits OpImageFetch instead), so the pipeline-creation failure is gone. The test then aborts on a Metal device limit, at the 3001-texel-wide volume case:
-[MTLTextureDescriptorInternal validateWithDevice:]:1421: failed assertion `Texture Descriptor Validation
MTLTextureDescriptor has width (3001) greater than the maximum allowed size of 2048.'
Metal caps a 3D texture at 2048 on this GPU (maxImageDimension3D is 2048 on MoltenVK here), and the 1/256-boundary case asks for 3001, which trips a hard Metal assertion that takes down the whole test. The fix itself is fine on MoltenVK — the gather lowering compiles and the pipeline builds; it is the test that can't run to completion here. A maxImageDimension3D guard on the wide-volume case (skip, or clamp the width) would let it run on Apple and on any device with a 2048 3D limit.
I could not get the test past that case to check the gathered values on MoltenVK, so this is not a value check — only the compile (red→green) and the device-limit abort.
GPUs tested: Apple M2 Max (MoltenVK 1.4.2).
Written with AI assistance (Claude Code). I ran the builds and tests above on my machine.
|
Thanks @jotitadev, the MoltenVK compile failure is a good fourth data point. You're right that 3001 was too wide: Vulkan only guarantees |
MarcZX10
left a comment
There was a problem hiding this comment.
The hardware evidence and the expanded coordinate/clamp coverage are unusually useful. I reviewed head a8779a382d3655ba1376b753cf76bc423247b56 and compiled the new focused test. I need two repository requirements completed before this is ready:
-
The checklist says this is based on current
mainata806c69bfe, but that is the merge base and the PR is 1,198 commits behind the currentmain(d6eafa54d0). The product build jobs for this head are skipped, so the reported device results do not validate the current merge base. Please update the branch, correct the checklist, and rerun the focused test against the updated code (or a current-main merge commit) on a supported host. My exact-head Release build ofagc_driver_image_gather_volume_testssucceeded. I attempted its CTest run with lavapipe, but the test process exits before Vulkan initialization because this host's Intel Xeon E5-1620 v2 lacks the required LZCNT, BMI1, BMI2, MOVBE, AVX2 and FMA extensions. I therefore could not independently execute the value checks here. -
CONTRIBUTING.md requires the Game bring-up discussion to list the open PRs and local changes the title still needs. The current PPSA02015 discussion (#5424) does not list this PR. Please update it with #9183 and the outstanding work needed to reach the shader; the PR description says the reproduction also needs #8223, #9186 and the local multisampled-target-retire change.
Local checks on this head: git diff --check passed; the conventions command exited 0 with one warning that the new debt entry is 1,026 characters, over the 1,000-character guideline. The focused test target compiled. Its execution was blocked by this host CPU check, as described above.
jotitadev
left a comment
There was a problem hiding this comment.
Re-ran a8779a3 on the M2 after the 241-wide change. The test now runs to the end on MoltenVK. The new wide case also fails without the exact product here, as it does on your GPUs.
Machine: Apple M2 Max, macOS 27.2, x86_64 under Rosetta 2. Release, BUILD_TESTING=ON.
Drivers: MoltenVK 1.4.2 (Vulkan SDK 1.4.363.0), and MoltenVK main + SPIRV-Cross #2704.
Base/method: PR #8658 feat/macos-port @ ee0a438415 (current main plus the macOS port), not main, which does not build on macOS. This PR's diff (git diff merge-base..a8779a382 -- . ':!docs') applied on top, clean. Full rebuild before each run.
- With the PR:
agc_driver_image_gather_volumepasses 3 of 3 on both drivers. The Metalwidth (3001) greater than the maximum allowed size of 2048abort from my earlier run is gone, so this is now a full value check on Apple, not just the compile. - Test only, production files at the merge base: fails on both drivers at pipeline creation with the same MSL error as before (
no member named 'gather' in 'metal::texture3d<...>',vkCreateComputePipelinesresult -3). - Recompiler from before the exact product (production files at c4381c7, test from a8779a3): the full test fails earlier, at
3D gather of red, wrap: thread 26 component 0 is 98/255, expected 2/255, because that commit also predates the addressing-mode reduction. So I ran only the 241-wide case through a local one-off patch (not part of the PR). It fails on both drivers with3D gather of red from a 241 texel wide volume at 1/256 texel boundaries: thread 0 component 0 is 243/255, expected 242/255. With the full PR, the same case alone passes 3 of 3. - Full
^agcsuite with the PR:-j1 --timeout 30, 857 tests. Base: 17 failures with MoltenVK 1.4.2 and 16 with main + #2704. With the PR there is one more on each driver:agc_driver_uint8x4_color_targeton 1.4.2 andagc_driver_cross_workgroup_flagon main + #2704. Both also fail intermittently on the base without this PR, with the same messages (a UINT16_ABGR export byte mismatch, anda workgroup did not see another workgroup's flag store). Neither involves a 3D gather. No timeouts.
GPUs tested: Apple M2 Max (MoltenVK 1.4.2; MoltenVK main + SPIRV-Cross #2704)
Written with AI assistance (Claude Code). I ran the builds and tests above on my machine.
a8779a3 to
2e3fb7f
Compare
2e3fb7f to
3f0d140
Compare
Dismissing a stale review: the requested changes were addressed after it (new head and/or updated description with the asked-for runs) and the merge queue blocks on it. Please request changes again if anything is still open.
|
Windows run on NVIDIA and the AMD proprietary driver, against current The merge onto current It needs a merge of
So on Windows both drivers fail without this change too: NVIDIA crashes as it does on Linux, and AMD's driver accepts the 3D AI-assisted: yes (Claude Code); results above were run on this machine. |
3f0d140 to
f2a7bda
Compare
|
Thanks @wachi35676, good catch. I rebased onto current |
|
Re-checked head
AI-assisted: yes (Claude Code); results above were run on this machine. |
f2a7bda to
0e27704
Compare
|
@fma-rc thanks for this one. We keep tests to real behaviour and argument validation, not asserting that a not-yet-implemented call throws, since that kind of test breaks the moment someone implements the function. In core/libs/prx/libSceAgcDriver/tests/execution/ImageGatherVolume.cpp please drop the two RunThrows checks (the 3D gather with a border color table, and the swizzled volume with border addressing) and the RunThrows helper that only they use, because both are still unimplemented paths in SpirvImageEmitter.cpp. The measured gather cases can stay. |
0e27704 to
15bb4cf
Compare
❌ NVIDIA GPU test: failedPR head
Failing testsThis PR changes test files, so each failing test was also built and run on main with only the PR's test changes applied. ✅ there means the failure comes from the PR's other changes; ❌ means the test fails on main too, so this run does not show the PR caused it.
agc_driver_image_gather_volume: Failed (also failed when re-run alone), assertion and diagnosticsguest_freetype: Subprocess aborted (also failed when re-run alone), assertion and diagnosticsUnstable on main (4, these also failed in recent main runs on this machine; not counted against this PR)Failing on main too (3, not counted against this PR)Validated but not compared (2, not in the main baseline's validation pass)Validation IDs already present on main (5 IDs in 717 tests)Skipped (51)ReproduceAutomated build and test on a contributor's NVIDIA machine; it does not replace upstream CI. A new comment is posted for each new push. Full logs. |
|
Hi fma-rc, A follow-up to the FreeType note. The NVIDIA GPU run at 345441c (merged onto
|
…snap their exact product
…nd keep z half border out of border handling A 3D sample of the same image through another sampler no longer feeds the gather's state, z modes 4 and 5 no longer count as border addressing, a null image no longer divides by a zero extent, and the test covers per-axis modes, rounding ties, denormals, null images and the vertex stage.
…every device's 3D image limit
…l descriptor field in the disk cache info encoder
345441c to
3cb6fb3
Compare
|
Thanks, @oneandonlydean, both were mine. The FreeType pointer came in through my rebase (a commit that picked up a submodule I had not updated); a new commit puts The wrap failure: the shader reduced a texel index of -1 with At 3cb6fb3 (rebased on |
XairRrR
left a comment
There was a problem hiding this comment.
Reviewed — approved
The 3D image_gather4_lz lowering merges cleanly, builds, and passes the focused image, gather and disk cache tests on lavapipe. The new test fails with main's recompiler. Two things were not checked here: the measured hardware model (no AMD GPU on this machine) and the bring-up list. The only new validation IDs come from main's sampler creation, not from this change.
What I ran
Head 3cb6fb3ebb3fa6d30b8a373c3a006504f6247445, main 4f41f4c375f5abb487dbd184fbc150f31994d762, merged tree 8d1e202a94108a3eba9754cfc8d4744b136fc146. Debian 13 VM (Xeon Gold 6136, AVX2/FMA), GCC 14.2.0, Release, Ninja. Vulkan was lavapipe (Mesa 26.1.6, LLVM 19); there is no GPU.
git diff --check 4f41f4c375f5abb487dbd184fbc150f31994d762...3cb6fb3ebb3fa6d30b8a373c3a006504f6247445
python3 tools/check_conventions.py --base 4f41f4c375f5abb487dbd184fbc150f31994d762
env VK_INSTANCE_LAYERS=VK_LAYER_KHRONOS_validation ANYPS5_NO_SHADER_CACHE=1 ANYPS5_REQUIRE_VULKAN=1 VK_LOADER_DEBUG=error,warn,layer ctest --test-dir build -V -R '^agc_driver_image_gather_volume$'
# negative control: the 8 production files restored from main
ninja -C build -j32 -l32 agc_driver_image_gather_volume_tests
ANYPS5_REQUIRE_VULKAN=1 ANYPS5_NO_SHADER_CACHE=1 ctest --test-dir build --output-on-failure -R '^agc_driver_image_gather_volume$'
# merged files restored, full rebuild
ninja -C build -j32 -l32
ANYPS5_REQUIRE_VULKAN=1 ctest --test-dir build -j8 --test-load 32 --output-on-failure -R 'gather|disk_cache|agc_driver_image_sample|agc_driver_image_table|shader_info|resource_tracker|descriptor_binding'
git status --porcelain
Results
| Check | Result |
|---|---|
git diff --check |
clean |
check_conventions.py |
All conventions checks passed |
Configure / build / libs (script) |
ok |
| Script test run (2163 tests) | no new failures, no timeouts; 7 known failures that also fail on main |
| Focused run (28 tests: gather, image sample, image table, disk cache) | 27 passed, 1 skipped (agc_driver_image_sample_lod_clamp, a device limitation) |
agc_driver_image_gather_volume |
passes in about 9.4 s on lavapipe with validation (limit is 30 s) |
| Validation layer | loaded; 8 new VUID- lines, all in the new test, all from vkCreateSampler with MIRROR_CLAMP_TO_EDGE (see Non-blocking) |
| Worktree after the control | git status --porcelain empty |
New validation IDs in agc_driver_image_gather_volume
Validation Error: [ VUID-VkSamplerCreateInfo-addressModeU-parameter ]
vkCreateSampler(): pCreateInfo->addressModeU (VK_SAMPLER_ADDRESS_MODE_MIRROR_CLAMP_TO_EDGE) requires the extensions VK_KHR_sampler_mirror_clamp_to_edge.
Validation Error: [ VUID-VkSamplerCreateInfo-addressModeV-parameter ] ...
Validation Error: [ VUID-VkSamplerCreateInfo-addressModeW-parameter ] ...
Validation Error: [ VUID-VkSamplerCreateInfo-addressModeU-01079 ]
vkCreateSampler(): pCreateInfo->addressModeU is VK_SAMPLER_ADDRESS_MODE_MIRROR_CLAMP_TO_EDGE but the VK_KHR_sampler_mirror_clamp_to_edge extension or samplerMirrorClampToEdge feature was not enabled.
Negative control
I restored the eight production files from main: Resources.hpp, ResourceMaterializer.hpp/.cpp, DescriptorBindingBuilder.cpp, ResourceTracker.cpp, PipelineSpecialization.hpp, ShaderDiskCache.cpp and SpirvImageEmitter.cpp. I kept the new test and rebuilt it. It fails on the first case:
Output
ShaderRecompiler::Recompile: SPIR-V validation before optimization failed:
level=2 word=607: [VUID-StandaloneSpirv-OpImage-04777] Expected Image 'Dim' to be 2D, Cube, or Rect
%531 = OpImageGather %v4float %529 %518 %uint_0
This matches the description's lavapipe result on main (pipeline refused). I then restored the merged files, rebuilt everything, and the test passed again.
Findings
Blocking
None.
Non-blocking
- The new test hits
main's missingsamplerMirrorClampToEdgeenablement (core/libs/prx/libSceAgcDriver/Graphics/src/GuestSamplerResource.cpp:32-44).- The test's mirror-once samplers (guest clamp modes 3, 5 and 7) go to
VK_SAMPLER_ADDRESS_MODE_MIRROR_CLAMP_TO_EDGEthroughtoVkAddressMode. - Nothing under
core/enablessamplerMirrorClampToEdgeorVK_KHR_sampler_mirror_clamp_to_edge. - The PR does not touch this code, and the gather itself never reads the sampler (it fetches). The 8 VUIDs therefore come from
mainand are listed as new only because the test does not exist there. - They still mean undefined behaviour in
vkCreateSamplerfor every guest sampler with those modes. Fixing it is a separate change: enable the feature, or throw when it is unavailable. It is worth a debt entry or a follow-up issue.
- The test's mirror-once samplers (guest clamp modes 3, 5 and 7) go to
- Unverified hardware claims are mentioned only in the debt entry. The debt entry in
docs/dev/debt/shader.mdcorrectly says the model comes from gfx1036, not the PS5's gfx1013. The 4.0 clamp bound and the half-border pinning are the parts most likely to differ between RDNA2 parts. A later run on a second RDNA2 GPU would firm them up. Not required for this PR.
Earlier review points, checked on this head:
- The
RunThrowschecks are gone. 3rdparty/freetypeis no longer in the diff.- The NVIDIA wrap failure from the RTX 2000 Ada run is addressed: the index reduction in
SpirvImageEmitter.cpp:925-928adds one period and usesOpUModwith a nonzero divisor (the extent isUMax(size, 1)). The test rejects anyOpSRemorOpSModin the gather shader.
I also checked by hand:
- The exact snap in
EmitVolumeGatherLz(SpirvImageEmitter.cpp:839): the rounding bias for negative values, the 64-bit shift split at shift 32, and denormals. - That the 11-bit gather state fits
RuntimeAbi::CompareStateMask(13 bits) for registration modules. - That the disk cache picks up the changed encoding through the generated source version.
Not verified
- Hardware model: I could not check the measured model (the 1/256 snap, half-border pinning, slice choice, border colors) against hardware: there is no AMD GPU here, and the evidence branch is outside the repository, so I did not rerun
checkall.py. Everything here is lavapipe only. - Other vendors: NVIDIA, RADV and MoltenVK results come from the author and other commenters, not from me.
- Bring-up list: I could not check whether the PPSA02015 bring-up discussion (#5424) now lists this PR, which an earlier review asked for (no network).
- In-game behaviour: not run (no titles are launched here).
✅ AMD GPU test: passedPR head
Unstable on main (2, these also failed in recent main runs on this machine; not counted against this PR)Failing on main too (8, not counted against this PR)Validated but not compared (1, not in the main baseline's validation pass)Validation IDs already present on main (4 IDs in 516 tests)Skipped (58)ReproduceAutomated build and test on a contributor's AMD machine; it does not replace upstream CI. A new comment is posted for each new push. Full logs. |
✅ NVIDIA GPU test: passedPR head
Unstable on main (4, these also failed in recent main runs on this machine; not counted against this PR)Failing on main too (3, not counted against this PR)Validated but not compared (1, not in the main baseline's validation pass)Validation IDs already present on main (4 IDs in 516 tests)Skipped (48)ReproduceAutomated build and test on a contributor's NVIDIA machine; it does not replace upstream CI. A new comment is posted for each new push. Full logs. |
What
image_gather4_lzthrough a 3D T# was emitted asOpImageGatheron a 3D image, which SPIR-V does not allow (gathers take 2D, cube and rect images). NVIDIA's compiler dies on it with SIGILL, RADV crashes, and lavapipe refuses the pipeline (vkCreateComputePipelines: Vulkan result -13). Death Stranding 2 (PPSA02015) has such a compute shader (0xc280112c00) and stops there: NVIDIA with SIGILL at 157 s, the 680M with a crash (rc 139).The hardware gathers four texels of one slice. Measured on a Radeon 610M (gfx1036) with the hardware oracle, 3742 rows over every clamp mode and border color:
The recompiler now emits four
OpImageFetchreads with that addressing. The sampler state reaches the shader as a specialization constant, or through the image's runtime metadata in registration modules. A border color table, a swizzled or integer texture with border addressing, image or sampler tables, unnormalized coordinates and samplers that disagree throw. The shader debt file records the measurement. Probes, results and the model check: evidence/gather-3d-volume (checkall.pymatches all 3742 rows).@mrfelek found the same shader in #5424 (comment) and fixed it on their fork (
fix/volume-gather-lz), without a PR, with four level-zero samples at the footprint's texel centres, from a Radeon 780M measurement of the slice choice and of wrap and clamp on it. This PR fetches the texels instead, because the measurement above shows that the sampler's filters, LOD clamp and bias do not change a gather, and it adds the per-axis modes, the 1/256 snap, the half-border modes and the border colors.Tested
Arch Linux, rebased on
main4f41f4c (2026-11-03, 3cb6fb3). Withmain's own newImageResourcefields the info encoder now counts 50 fields and the record stays 108 bytes. The two checks that expected the unimplemented border paths to throw are dropped. The FreeType submodule ismain's again (my earlier rebase had moved it back to 2.13.3).agc_driver_image_*,agc_shader_*,agc_shader_disk_cache*) pass on the RTX 5080, RADV (Radeon 610M) and lavapipe, except three that skip on RADV (agc_driver_image_array_rows_past_layer_limit,agc_driver_image_array_window_past_layer_limit,agc_driver_image_table_vouched_lookups). The Windows build and ctest pass (3919 of 3919), with no failure the merge base lacks.OpSRemwith a negative operand is undefined in Vulkan withoutmaintenance8, and the NVIDIA RTX 2000 Ada in CI returned 3 for -1 mod 7.agc_driver_image_gather_volumenow also fails if the gather shader containsOpSRemorOpSMod; it fails at e92ec41 (test only) and passes with the fix.agc_driver_image_gather_volume: red channel, alpha and blue gathers over all eight clamp modes and three border colors, an odd 7x5x3 volume, per-axis modes, rounding ties and denormal coordinates, a 241-texel-wide volume at 1/256 boundaries, a base level, a swizzle, null images, a vertex-shader variant through the registration path, and a gather next to a 3D sample through another sampler. Expected values come from the measured model.mainthe test fails on every device: NVIDIA RTX 5080 SIGILL (132), Radeon 610M and 680M crash (139), lavapipevkCreateComputePipelines: Vulkan result -13, llvmpipe rc 1. With the PR it passes on all five.main's extent and passes with the fix.agc_driver_image_gather_one_dimensional,image_converted_unorm,image_sample_reduction,image_sint8,image_srgb_load,image_sample_unnormalizedpass onmainand with the PR on the RTX 5080, 610M and lavapipe.conventions_env_reference,conventions_blob_reads,agc_driver_f64_denorm_preserve_spirv,agc_driver_f16_to_f32_unpack_spirv, the threenet_dontwait_blocking_send*) fail the same way on the merge base on this machine.mainwith PR 5426, PR 7978, PR 6565, PR 6182, PR 8223 and our multisampled-target retire change, which it needs to reach this shader: without this PR the RTX 5080 stops with SIGILL in the shader compiler at 157 s and the 680M crashes; with it both run on to the next stop,color targets with fewer fragments than samples (EQAA) are unsupported(RTX 5080 at 86 s, frame 88; 680M at 228 s, frame 63, picture visible).tools/check_conventions.py --base origin/mainpasses.Checklist
main; no other open PR implements the same functionsNotImplemented_nid_no_patch); silent stubs are listed under Silent stubs in their library's technical debt file.mdfiles or images added to the repository (attach them to this PR)